#ifndef OCL_H
#define OCL_H

#include <assert.h>

#include <algorithm>
#include <iostream>
#include <string>
#include <vector>
#include <map>

#ifdef USE_OCL_SL
    #include "ocl_sl.h"
#else
    #include "ocl_dl.h"
#endif

#include "ocl_comm.h"
#include "ocl_info.h"
#include "ocl_type.h"

namespace ocl
{
//=======================================
//  oclManager
//      |
//      |--m_selDev : oclDevice
//                        |
//                        |--m_context : cl_context
//                        |--m_cmdque  : cl_command_queue
//
//  oclProgram
//      |
//      |--m_refDev : oclDevice&
//      |--m_progrm : cl_program
//      |--m_mapKrl : std::map<std::string, oclKernel*>
//                        |
//                        |-- oclKernel
//                               |
//                               |--m_refDev : oclDevice&
//                               |--m_kernel : cl_kernel
//  oclBuf
//      |
//      |--R()
//      |--W()
//      |--G()
//      |--CP() : Read buffer from HW_mem into SYS_mem
//  oclImg
//      |
//      |--G()
//      |--CP() : Read buffer from HW_mem into SYS_mem

//=======================================
//  oclDevice
//=======================================
class oclDevice
{
public:
    oclDevice(cl_platform_id pID, cl_device_id dID)
        : m_pfmID(NULL), m_devID(NULL), m_context(NULL), m_cmdque(NULL)
    {
        if (pID == NULL || dID == NULL)
            return;

        if (m_pfmID != NULL || m_devID != NULL)
            Clean();

        m_pfmID = pID;
        m_devID = dID;
        m_info.Init(m_devID);
    };
    ~oclDevice() {Clean();};
    void Setup()
    {
        assert(m_context == NULL && m_cmdque == NULL);
        cl_int err = 0;
        m_context = oclCreateContext(NULL, 1, &m_devID, NULL, NULL, &err);
        ocl::logErr(err);
        m_cmdque = oclCreateCommandQueue(m_context, m_devID, 0/*CL_QUEUE_PROFILING_ENABLE*/, &err);
        ocl::logErr(err);
    };
    void Clean()
    {
        if (m_cmdque)   OCL_CHKRET(oclReleaseCommandQueue(m_cmdque));
        if (m_context)  OCL_CHKRET(oclReleaseContext(m_context));
        m_pfmID = NULL;
        m_devID = NULL;
    };
    bool			 Ready()				{return (m_context != NULL && m_cmdque != NULL);};
    oclDeviceInfo&   GetInfo()              {return m_info;};
    cl_platform_id   Get_cl_platform_id()   {return m_pfmID;};
    cl_device_id     Get_cl_device_id()     {return m_devID;};
    cl_context       Get_cl_context()		{return m_context;};
    cl_command_queue Get_cl_command_queue() {return m_cmdque;};

    bool operator==(oclDevice& rhs)
    {
        return GetInfo() == rhs.GetInfo();
    };
    bool operator!=(oclDevice& rhs)
    {
        return !(*this == rhs);
    };

private:
    oclDevice(const oclDevice& rhs);
    oclDevice& operator=(const oclDevice& rhs);

    cl_platform_id      m_pfmID;
    cl_device_id        m_devID;  
    cl_context          m_context;
    cl_command_queue    m_cmdque;
    oclDeviceInfo       m_info;
};
//=======================================
//  oclKernel
//=======================================
class oclKernel
{
public:
    oclKernel(cl_kernel clkerl, oclDevice& oclDev): m_kernel(clkerl), m_refDev(oclDev), m_dim(0)
    {
        if (!m_kernel)
            return;

        OCL_CHKRET(oclRetainKernel(m_kernel));
        m_info.Init(m_kernel, m_refDev.Get_cl_device_id());
    };
    ~oclKernel()
    {
        if (m_kernel)
            OCL_CHKRET(oclReleaseKernel(m_kernel));
    };
    void SetDimensionAndSize(unsigned int dim, size_t* pGlobal, size_t* pLocal)
    {
        m_dim = dim;
        m_global.clear();
        m_local.clear();
        for(unsigned int i = 0 ; i < dim ; ++i)
        {
            if (pGlobal[i] <= 0)
                throw std::runtime_error("pGlobal[i]<=0");
            if (pLocal[i] <= 0)
                throw std::runtime_error("pLocal[i]<=0");

            m_global.push_back(pGlobal[i]);
            m_local .push_back(pLocal[i]);
        }
    };
	template<typename T>
	void SetArg(unsigned int idx ,T* arg)
	{
		if (idx >= m_info.oCL_KERNEL_NUM_ARGS)
			throw std::runtime_error("Argument index exceed the kernel function range!");
		OCL_CHKRET(oclSetKernelArg(m_kernel, idx, sizeof(T), (void *)arg));
	};
	void Run(bool waitForFinish = false)
    {
        assert(m_dim != 0);
        cl_command_queue cmdque = m_refDev.Get_cl_command_queue();
        cl_event evt = 0;
        cl_int   err = 0;
        // Let vendor to determine best local size in GPU driver level.
        err = oclEnqueueNDRangeKernel(cmdque, m_kernel, m_dim, 0, &m_global[0], 0/*&m_local[0]*/, 0, 0, waitForFinish?&evt:0);
        ocl::logErr(err);
        if (evt)
        {
            oclWaitForEvents(1, &evt);
            oclReleaseEvent(evt);
        }
    };

private:
    oclKernel(const oclKernel&);
    oclKernel& operator=(const oclKernel&);

    cl_kernel        m_kernel;
    oclDevice&       m_refDev;
    oclKernelInfo    m_info;
    unsigned int          m_dim;
    std::vector<size_t>   m_global;
    std::vector<size_t>   m_local;
};
//=======================================
//  oclProgram
//=======================================
class oclProgram
{
public:
    oclProgram(oclDevice& dev, const char* progStr, const char* options = NULL): m_progrm(NULL), m_refDev(dev)
    {
        cl_int err = 0;
        m_progrm = oclCreateProgramWithSource(dev.Get_cl_context(), 1, &progStr, 0, &err);
        ocl::logErr(err);

        if((err = oclBuildProgram(m_progrm, 0, 0, options, 0, 0)) != CL_SUCCESS)
        {
            size_t	errlen;
            char	errbuf[2048];
            oclGetProgramBuildInfo(m_progrm, m_refDev.Get_cl_device_id(), 
                CL_PROGRAM_BUILD_LOG, sizeof(errbuf), errbuf, &errlen);
            ocl::dprintf("clGetProgramBuildInfo=%s\n", errbuf);
            throw std::runtime_error("[ERROR] clBuildProgram FAILED\n");
        }

        cl_uint kerlCnt = 0;
        std::vector<cl_kernel> kernels;
        OCL_CHKRET(oclCreateKernelsInProgram(m_progrm, 0, 0, &kerlCnt));
        if (kerlCnt)
        {
            kernels.resize(kerlCnt);
            OCL_CHKRET(oclCreateKernelsInProgram(m_progrm, kerlCnt, &kernels[0], 0));
            for(unsigned int i = 0; i < kernels.size(); ++i)
            {
                char kerlName[OCL_NAME_MAX] = {0};
                OCL_CHKRET(oclGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, OCL_NAME_MAX, kerlName, 0));
                m_mapKrl[kerlName] = new oclKernel(kernels[i], m_refDev);
            }
        }

        std::vector<cl_kernel>::const_iterator cIt = kernels.begin();
        for(; cIt != kernels.end(); ++cIt)
            oclReleaseKernel(*cIt);

        // sentinal NULL_KERNEL
        m_mapKrl["OCL_NULL_KERNEL"] = new oclKernel(NULL, m_refDev);
    };
    ~oclProgram()
    {
        std::map<std::string, oclKernel*>::iterator it;
        for ( it = m_mapKrl.begin(); it != m_mapKrl.end(); ++it )
            delete it->second;
        m_mapKrl.clear();

        if (m_progrm)
        {
            OCL_CHKRET(oclReleaseProgram(m_progrm));
            m_progrm = NULL;
        }
    };
    oclKernel& operator[](const char* kerlName)
    {
        if (m_mapKrl.find(kerlName) != m_mapKrl.end())
        {
            return *m_mapKrl[kerlName];
        }
        return *m_mapKrl["OCL_NULL_KERNEL"];
    };
    oclDevice& GetDevice() const { return m_refDev;};
private:
    oclProgram(const oclProgram&);
    oclProgram& operator=(const oclProgram&);

    cl_program                        m_progrm;
    oclDevice&                        m_refDev;
    std::map<std::string, oclKernel*> m_mapKrl;
};
//=======================================
//  oclManager
//=======================================
class oclManager
{
public:
    static oclManager& Instance()
    {
        static oclManager mgr;
        return mgr;
    };
    oclDevice& SelDevice(const cl_device_type devType, unsigned int seqOrder = 0)
    {
        unsigned int totalDev = m_vecDev.size();

        if (totalDev == 0)
            throw std::runtime_error("[ERROR] No available OpenCL device.");
        if (devType == CL_DEVICE_TYPE_DEFAULT)
            throw std::runtime_error("[ERROR] OCL only support CL_DEVICE_TYPE_[CPU | GPU | ALL].");
        if (devType == CL_DEVICE_TYPE_CPU && m_nCPUDev == 0)
            throw std::runtime_error("[ERROR] No available CL_DEVICE_TYPE_CPU device.");
        if (devType == CL_DEVICE_TYPE_GPU && m_nGPUDev == 0)
            throw std::runtime_error("[ERROR] No available CL_DEVICE_TYPE_GPU device.");

        if (devType == CL_DEVICE_TYPE_CPU && seqOrder >= m_nCPUDev)
            seqOrder %= m_nCPUDev;
        else if (devType == CL_DEVICE_TYPE_GPU && seqOrder >= m_nGPUDev)
            seqOrder %= m_nGPUDev;
        else if (devType == CL_DEVICE_TYPE_ALL && seqOrder >= totalDev )
            seqOrder %= totalDev;

        unsigned int avalCnt = 0;
        for( unsigned int idx = 0; idx < m_vecDev.size(); ++idx)
        {
            oclDeviceInfo& info = m_vecDev[idx]->GetInfo();
            if (info.oCL_DEVICE_TYPE & devType)
            {
                ++avalCnt;
                if (seqOrder == avalCnt-1)
                {
                    if (!m_vecDev[idx]->Ready())
                        m_vecDev[idx]->Setup();
                    m_selDev = m_vecDev[idx];
                    break;
                }
            }
        }
        return *m_selDev;
    };

    inline oclDevice&   GetDevice()   {return *m_selDev;};
    inline void         GetDevInfo(int idx, oclDeviceInfo& info)   {info = m_vecDev[idx]->GetInfo();};
    inline void         GetPfmInfo(int idx, oclPlatformInfo& info) {info = m_vecPfm[idx];};
    inline unsigned int GetPfmCount() {return m_vecPfm.size();};
    inline unsigned int GetDevCount() {return m_vecDev.size();};

private:
    oclManager():m_selDev(NULL),m_nGPUDev(0),m_nCPUDev(0){setup();};
    ~oclManager()                {clean();};
    oclManager(const oclManager&);
    oclManager& operator=(const oclManager&);
    
    void setup()
    {
        cl_int err = 0;
        cl_uint pfmCnt = 0;

        err = oclGetPlatformIDs(0, 0, &pfmCnt);
        if(err != CL_SUCCESS) {
            std::cerr << "ERROR: Calc pfmCnt by oclGetPlatformIDs()\n";
            return;
        }
        // no available platform found
        if (pfmCnt <= 0)
            return;

        std::vector<cl_platform_id> vPfmId(pfmCnt);
        m_vecPfm.resize(pfmCnt);
        err = oclGetPlatformIDs(pfmCnt, &vPfmId[0], &pfmCnt);
        if(err != CL_SUCCESS) {
            std::cerr << "ERROR:oclGetPlatformIDs()\n";
            return;
        }
        
        for(cl_uint i = 0; i < pfmCnt; i++)
        {
            m_vecPfm[i].Init(vPfmId[i]);

            unsigned int devCnt = 0;
            OCL_CHKRET(oclGetDeviceIDs(vPfmId[i], CL_DEVICE_TYPE_ALL, 0, 0, &devCnt));

            // no available device found
            if (devCnt <= 0)
                continue;

            std::vector<cl_device_id> vDev(devCnt);
            OCL_CHKRET(oclGetDeviceIDs(vPfmId[i], CL_DEVICE_TYPE_ALL, devCnt, &vDev[0], &devCnt));
            for(unsigned int devIdx = 0; devIdx < devCnt; ++devIdx)
            {
                oclDevice* tmp = new oclDevice(vPfmId[i], vDev[devIdx]);
                if (tmp->GetInfo().oCL_DEVICE_TYPE & CL_DEVICE_TYPE_GPU) m_nGPUDev++;
                if (tmp->GetInfo().oCL_DEVICE_TYPE & CL_DEVICE_TYPE_CPU) m_nCPUDev++;
                m_vecDev.push_back(tmp);
            }
        }
    };
    void clean()
    {
        for(unsigned int i = 0 ; i < m_vecDev.size(); ++i)
            delete m_vecDev[i];
        m_vecDev.clear();
    };

    
    std::vector<oclPlatformInfo>    m_vecPfm;
    std::vector<oclDevice*>         m_vecDev;
    oclDevice*                      m_selDev;
    unsigned int                    m_nGPUDev;
    unsigned int                    m_nCPUDev;
};
//=======================================
//  oclMem
//=======================================
struct oclMem
{
    static void unmap(cl_mem src, void* dst)
    {
        cl_command_queue cmdque = oclManager::Instance().GetDevice().Get_cl_command_queue();
        if (!cmdque)
            return;

        OCL_CHKRET(oclEnqueueUnmapMemObject(cmdque, src, dst, 0, 0, 0));
    };
};
//=======================================
//  oclBuf
//=======================================
struct oclBuf: public oclMem
{
    // Specialized for creating: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR
    template<typename T>
    static cl_mem R(T* buf, unsigned int bufLen)
    {
        cl_context ctx = oclManager::Instance().GetDevice().Get_cl_context();
        if (!ctx)
            return NULL;

        cl_int err = 0;
        cl_mem mem = oclCreateBuffer(ctx,
                                    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                    sizeof(typename Std2Ocl<T>::ResultType) * bufLen,
                                    buf,
                                    &err);
        ocl::logErr(err);
        return mem;
    };
    // Specialized for creating: CL_MEM_WRITE_ONLY
    template<typename T>
    static cl_mem W(T* /*buf*/, unsigned int bufLen)
    {
        cl_context ctx = oclManager::Instance().GetDevice().Get_cl_context();
        if (!ctx)
            return NULL;

        cl_int err = 0;
        cl_mem mem = oclCreateBuffer(ctx,
                                    CL_MEM_WRITE_ONLY,
                                    sizeof(typename Std2Ocl<T>::ResultType) * bufLen,
                                    NULL,
                                    &err);
        ocl::logErr(err);
        return mem;
    };
    template<typename T>
    static cl_mem create(T* buf, unsigned int bufLen, int flags)
    {
        cl_context ctx = oclManager::Instance().GetDevice().Get_cl_context();
        if (!ctx)
            return NULL;

        cl_int err = 0;
        cl_mem mem = oclCreateBuffer(ctx,
                                    flags,
                                    sizeof(typename Std2Ocl<T>::ResultType) * bufLen,
                                    buf,
                                    &err);
        ocl::logErr(err);
        return mem;
    };
    template<typename T>
    static void read(cl_mem src, T* dst, unsigned int len, unsigned int off = 0)
    {
        cl_command_queue cmdque = oclManager::Instance().GetDevice().Get_cl_command_queue();
        if (!cmdque)
            return;
        OCL_CHKRET(oclEnqueueReadBuffer(cmdque, src, CL_TRUE, sizeof(T)*off, sizeof(T)*len, dst, 0, 0, 0));
    };
    static void map(cl_mem src, void*& dst,
                    cl_map_flags flags,
                    size_t offset, size_t lenInBytes)
    {
        cl_command_queue cmdque = oclManager::Instance().GetDevice().Get_cl_command_queue();
        if (!cmdque)
            return;

        cl_int err = 0;
        dst = oclEnqueueMapBuffer(cmdque, src, CL_TRUE, flags, offset, lenInBytes, 0, 0, 0, &err);
        ocl::logErr(err);
    };
    template<typename T>
    static int cmp(cl_mem mem, T* tar, unsigned int len, T tolerance = 0, unsigned int off = 0)
    {
        if (tar == NULL)
            return -1;
        T* src = new (std::nothrow) T[len];
        if (src == NULL)
            return -1;

        oclBuf::read(mem, src, len, off);

        unsigned int retIdx = 0;
        for(unsigned int i = 0; i < len; ++i)
        {
            T d = src[i] - tar[i];
            if (d > tolerance || d < -tolerance)
            {
                retIdx = i;
                ocl::dprintf("[%d]\n", i);
                break;
            }
        }
        delete[] src;
        return retIdx;
    };
};
//=======================================
//  oclImg
//=======================================
struct oclImg: public oclMem
{
    static cl_mem create(void* ImgPtr,
                         unsigned int ImgW,
                         unsigned int ImgH,
                         unsigned int ImgS,
                         int flags,
                         cl_image_format* ImgFmt)
    {
        cl_context ctx = oclManager::Instance().GetDevice().Get_cl_context();
        if (!ctx)
            return NULL;

        cl_int err = 0;
        cl_mem mem = oclCreateImage2D(ctx, flags, ImgFmt, ImgW, ImgH, ImgS, ImgPtr, &err);
        ocl::logErr(err);
        return mem;
    };
    static void read(cl_mem src,
                     void*  dst,
                     size_t x, size_t y,
                     size_t w, size_t h)
    {
        cl_command_queue cmdque = oclManager::Instance().GetDevice().Get_cl_command_queue();
        if (!cmdque)
            return;

        size_t origin[3] = {x, y, 0};
        size_t region[3] = {w, h, 1};
        OCL_CHKRET(oclEnqueueReadImage(cmdque, src, CL_TRUE, origin, region, 0, 0, dst, 0, 0, 0));
    };
    static void map(cl_mem src, void*& dst,
                    cl_map_flags flags,
                    size_t x, size_t y,
                    size_t w, size_t h,
                    size_t* r = NULL, size_t* s = NULL)
    {
        cl_command_queue cmdque = oclManager::Instance().GetDevice().Get_cl_command_queue();
        if (!cmdque)
            return;

        cl_int err = 0;
        size_t origin[3] = {x, y, 0};
        size_t region[3] = {w, h, 1};
        size_t row_pitch = 0;
        size_t slice_pitch = 0;
        dst = oclEnqueueMapImage(cmdque, src, CL_TRUE, flags, origin, region, &row_pitch, &slice_pitch, 0, 0, 0, &err);
        if (r)    *r = row_pitch;
        if (s)    *s = slice_pitch;
        ocl::logErr(err);
    };
};

oclManager&			oclGetManager();
oclDevice&			oclGetDevice();
oclDevice&			oclSelDevice(cl_device_type devType, unsigned int seqOrder = 0);
cl_context			oclGetContext();
cl_command_queue	oclGetCommandQueue();

void                oclGetPfmInfo(unsigned int idx, oclPlatformInfo& info);
void                oclGetDevInfo(unsigned int idx, oclDeviceInfo&   info);
unsigned int        oclGetPfmCount();
unsigned int        oclGetDevCount();
}
#endif
